<!-- HTML header for doxygen 1.8.13-->
<!DOCTYPE html PUBLIC "-//W3C//DTD XHTML 1.0 Transitional//EN" "http://www.w3.org/TR/xhtml1/DTD/xhtml1-transitional.dtd">
<html xmlns="http://www.w3.org/1999/xhtml">
<head>
<meta http-equiv="Content-Type" content="text/xhtml;charset=UTF-8"/>
<meta http-equiv="X-UA-Compatible" content="IE=9"/>
<meta name="generator" content="Doxygen 1.8.14"/>
<meta name="viewport" content="width=device-width, initial-scale=1"/>
<title>Taskflow Handbook</title>
<link href="tabs.css" rel="stylesheet" type="text/css"/>
<link rel="icon" type="image/x-icon" href="favicon.ico" />
<script type="text/javascript" src="jquery.js"></script>
<script type="text/javascript" src="dynsections.js"></script>
<link href="navtree.css" rel="stylesheet" type="text/css"/>
<script type="text/javascript" src="resize.js"></script>
<script type="text/javascript" src="navtreedata.js"></script>
<script type="text/javascript" src="navtree.js"></script>
<script type="text/javascript">
/* @license magnet:?xt=urn:btih:cf05388f2679ee054f2beb29a391d25f4e673ac3&amp;dn=gpl-2.0.txt GPL-v2 */
  $(document).ready(initResizable);
/* @license-end */</script>
<link href="search/search.css" rel="stylesheet" type="text/css"/>
<script type="text/javascript" src="search/searchdata.js"></script>
<script type="text/javascript" src="search/search.js"></script>
<link href="doxygen.css" rel="stylesheet" type="text/css" />
</head>
<body>
<div id="top"><!-- do not remove this div, it is closed by doxygen! -->
<div id="titlearea">
<table cellspacing="0" cellpadding="0">
 <tbody>
 <tr style="height: 56px;">
  <td id="projectalign" style="padding-left: 0.5em;">
   <div id="projectname"><a href="https://taskflow.github.io/">Taskflow</a>
   &#160;<span id="projectnumber">3.0.0-Master-Branch</span>
   </div>
  </td>
 </tr>
 </tbody>
</table>
</div>
<!-- end header part -->
<!-- Generated by Doxygen 1.8.14 -->
<script type="text/javascript">
/* @license magnet:?xt=urn:btih:cf05388f2679ee054f2beb29a391d25f4e673ac3&amp;dn=gpl-2.0.txt GPL-v2 */
var searchBox = new SearchBox("searchBox", "search",false,'Search');
/* @license-end */
</script>
<script type="text/javascript" src="menudata.js"></script>
<script type="text/javascript" src="menu.js"></script>
<script type="text/javascript">
/* @license magnet:?xt=urn:btih:cf05388f2679ee054f2beb29a391d25f4e673ac3&amp;dn=gpl-2.0.txt GPL-v2 */
$(function() {
  initMenu('',true,false,'search.php','Search');
  $(document).ready(function() { init_search(); });
});
/* @license-end */</script>
<div id="main-nav"></div>
</div><!-- top -->
<div id="side-nav" class="ui-resizable side-nav-resizable">
  <div id="nav-tree">
    <div id="nav-tree-contents">
      <div id="nav-sync" class="sync"></div>
    </div>
  </div>
  <div id="splitbar" style="-moz-user-select:none;" 
       class="ui-resizable-handle">
  </div>
</div>
<script type="text/javascript">
/* @license magnet:?xt=urn:btih:cf05388f2679ee054f2beb29a391d25f4e673ac3&amp;dn=gpl-2.0.txt GPL-v2 */
$(document).ready(function(){initNavTree('chapter7.html','');});
/* @license-end */
</script>
<div id="doc-content">
<!-- window showing the filter options -->
<div id="MSearchSelectWindow"
     onmouseover="return searchBox.OnSearchSelectShow()"
     onmouseout="return searchBox.OnSearchSelectHide()"
     onkeydown="return searchBox.OnSearchSelectKey(event)">
</div>

<!-- iframe showing the search results (closed by default) -->
<div id="MSearchResultsWindow">
<iframe src="javascript:void(0)" frameborder="0" 
        name="MSearchResults" id="MSearchResults">
</iframe>
</div>

<div class="header">
  <div class="headertitle">
<div class="title">C7: GPU Tasking (cudaFlowCapturer) </div>  </div>
</div><!--header-->
<div class="contents">
<div class="textblock"><p>You can create a cudaFlow through <em>stream capture</em>, which allows you to capture information on GPU activities that are submitted to the stream managed by a <em>cudaFlowCapturer</em>.</p>
<h1><a class="anchor" id="C7_Capture_a_cudaFlow"></a>
Capture a cudaFlow</h1>
<p>When your program has no access to direct kernel calls but invoke it through a stream-based interface (e.g., <a href="https://docs.nvidia.com/cuda/cublas/index.html">cuBLAS</a> and <a href="https://developer.nvidia.com/cudnn">cuDNN</a> library functions), you can use <a class="el" href="classtf_1_1cudaFlowCapturer.html" title="class for building a CUDA task dependency graph through stream capture ">tf::cudaFlowCapturer</a> to capture the GPU activities into a cudaFlow. A cudaFlowCapturer is similar to a cudaFlow except it forms a GPU task graph through <em>stream capture</em>. You use the method <a class="el" href="classtf_1_1cudaFlowCapturerBase.html#adf651356def71f613c589c29588398c2" title="captures a sequential CUDA operations from the given callable ">tf::cudaFlowCapturer::on</a> to capture a sequence of <em>asynchronous</em> CUDA operations through the given stream.</p>
<p>The following example creates a CUDA graph that captures two kernel tasks, <code>task_1</code> and <code>task_2</code>, where <code>task_1</code> (i.e., <code>my_kernel_1</code>) runs before <code>task_2</code> (i.e., <code>my_kernel_2</code>).</p>
<div class="fragment"><div class="line"><a class="code" href="classtf_1_1Task.html">tf::Task</a> task = taskflow.<a class="code" href="classtf_1_1FlowBuilder.html#a60d7a666cab71ecfa3010b2efb0d6b57">emplace</a>([&amp;](<a class="code" href="classtf_1_1cudaFlowCapturer.html">tf::cudaFlowCapturer</a>&amp; capturer){</div><div class="line">  <span class="comment">// capture my_kernel_1</span></div><div class="line">  <a class="code" href="classtf_1_1cudaTask.html">tf::cudaTask</a> task_1 = capturer.<a class="code" href="classtf_1_1cudaFlowCapturerBase.html#adf651356def71f613c589c29588398c2">on</a>([&amp;](cudaStream_t stream){  <span class="comment">// stream is managed by the capturer</span></div><div class="line">    my_kernel_1&lt;&lt;&lt;grid_1, block_1, shm_size_1, stream&gt;&gt;&gt;(my_parameters_1);</div><div class="line">  });</div><div class="line"></div><div class="line">  <span class="comment">// capture my_kernel_2</span></div><div class="line">  <a class="code" href="classtf_1_1cudaTask.html">tf::cudaTask</a> task_2 = capturer.<a class="code" href="classtf_1_1cudaFlowCapturerBase.html#adf651356def71f613c589c29588398c2">on</a>([&amp;](cudaStream_t stream){  <span class="comment">// stream is managed by the capturer</span></div><div class="line">    my_kernel_2&lt;&lt;&lt;grid_2, block_2, shm_size_2, stream&gt;&gt;&gt;(my_parameters_2);</div><div class="line">  });</div><div class="line"></div><div class="line">  <span class="comment">// my_kernel_1 runs before my_kernel_2</span></div><div class="line">  task_1.<a class="code" href="classtf_1_1cudaTask.html#abdd68287ec4dff4216af34d1db44d1b4">precede</a>(task_2);</div><div class="line">});</div></div><!-- fragment --><p>Inside <a class="el" href="classtf_1_1cudaFlowCapturerBase.html#adf651356def71f613c589c29588398c2" title="captures a sequential CUDA operations from the given callable ">tf::cudaFlowCapturer::on</a>, you should <em>NOT</em> modify the properties of the stream argument but use it to capture <em>asynchronous</em> GPU operations (e.g., <code>kernel</code>, <code>cudaMemcpyAsync</code>). The stream object passed to each <a class="el" href="classtf_1_1cudaFlowCapturerBase.html#adf651356def71f613c589c29588398c2" title="captures a sequential CUDA operations from the given callable ">tf::cudaFlowCapturer::on</a> call may differ, and it depends on how the internal optimization algorithm maximizes the GPU parallelism.</p>
<p>A cudaFlowCapturer lives with the callable. When the executor invoke the capturer callable, it creates the cudaFlowCapturer and will destroy it until all internal operations finish.</p>
<h1><a class="anchor" id="C7_CaptureWithinAcudaFlow"></a>
Capture within a cudaFlow</h1>
<p>Within a parent cudaFlow, you can capture a cudaFlow to form a subflow that eventually becomes a <em>child</em> node in the underlying CUDA task graph. The following example defines a captured flow <code>task2</code> of two dependent tasks, <code>task2_1</code> and <code>task2_2</code>, and <code>task2</code> runs after <code>task1</code>.</p>
<div class="fragment"><div class="line"><a class="code" href="classtf_1_1Task.html">tf::Task</a> task = taskflow.<a class="code" href="classtf_1_1FlowBuilder.html#a60d7a666cab71ecfa3010b2efb0d6b57">emplace</a>([&amp;](<a class="code" href="classtf_1_1cudaFlow.html">tf::cudaFlow</a>&amp; cf){</div><div class="line"></div><div class="line">  <a class="code" href="classtf_1_1cudaTask.html">tf::cudaTask</a> task1 = cf.<a class="code" href="classtf_1_1cudaFlow.html#adb731be71bdd436dfb5e36e6213a9a17">kernel</a>(grid, block, shm, my_kernel, args...);</div><div class="line"></div><div class="line">  <a class="code" href="classtf_1_1cudaTask.html">tf::cudaTask</a> task2 = cf.<a class="code" href="classtf_1_1cudaFlow.html#a89c389fff64a16e5dd8c60875d3b514d">capture</a>([&amp;](<a class="code" href="classtf_1_1cudaFlowCapturer.html">tf::cudaFlowCapturer</a>&amp; capturer){</div><div class="line">    <a class="code" href="classtf_1_1cudaTask.html">tf::cudaTask</a> task2_1 = capturer.<a class="code" href="classtf_1_1cudaFlowCapturerBase.html#adf651356def71f613c589c29588398c2">on</a>([&amp;](cudaStream_t stream){  </div><div class="line">      my_kernel2&lt;&lt;&lt;grid1, block1, shm_size1, stream&gt;&gt;&gt;(args1...);</div><div class="line">    });</div><div class="line">    <a class="code" href="classtf_1_1cudaTask.html">tf::cudaTask</a> task2_2 = capturer.<a class="code" href="classtf_1_1cudaFlowCapturerBase.html#adf651356def71f613c589c29588398c2">on</a>([&amp;](cudaStream_t stream){  </div><div class="line">      my_kernel2&lt;&lt;&lt;grid2, block2, shm_size2, stream&gt;&gt;&gt;(args2...);</div><div class="line">    });</div><div class="line">    task2_1.<a class="code" href="classtf_1_1cudaTask.html#abdd68287ec4dff4216af34d1db44d1b4">precede</a>(task2_2);</div><div class="line">  });</div><div class="line"></div><div class="line">  task1.<a class="code" href="classtf_1_1cudaTask.html#abdd68287ec4dff4216af34d1db44d1b4">precede</a>(task2);</div><div class="line">});</div></div><!-- fragment --><h1><a class="anchor" id="C7_CommonCaptureMethods"></a>
Common Capture Methods</h1>
<p>A cudaFlowCapturer inherits <a class="el" href="classtf_1_1cudaFlowCapturerBase.html" title="base class to construct a CUDA task graph through stream capture ">tf::cudaFlowCapturerBase</a> that defines a set of methods for capturing common GPU operations, including <a class="el" href="classtf_1_1cudaFlowCapturerBase.html#ad836d32cc2e9532ce57fe3ad6cc67d5d" title="captures a kernel ">tf::cudaFlowCapturerBase::kernel</a>, <a class="el" href="classtf_1_1cudaFlowCapturerBase.html#a38adf66dfcc0829708db653d153a83e2" title="copies data between host and device asynchronously through a stream ">tf::cudaFlowCapturerBase::memcpy</a>, and <a class="el" href="classtf_1_1cudaFlowCapturerBase.html#a8e3071171c0875c93dcc077a2e0a435a" title="initializes or sets GPU memory to the given value byte by byte ">tf::cudaFlowCapturerBase::memset</a>. For example, the following code snippet constructs a GPU task graph of one host-to-device copy, kernel, and one device-to-host copy, in this order of their dependencies.</p>
<div class="fragment"><div class="line"><a class="code" href="classtf_1_1Task.html">tf::Task</a> task = taskflow.<a class="code" href="classtf_1_1FlowBuilder.html#a60d7a666cab71ecfa3010b2efb0d6b57">emplace</a>([](<a class="code" href="classtf_1_1cudaFlowCapturer.html">tf::cudaFlowCapturer</a>&amp; capturer){</div><div class="line">  <span class="comment">// copy data from host_data to gpu_data</span></div><div class="line">  <a class="code" href="classtf_1_1cudaTask.html">tf::cudaTask</a> h2d = capturer.<a class="code" href="classtf_1_1cudaFlowCapturerBase.html#a38adf66dfcc0829708db653d153a83e2">memcpy</a>(gpu_data, host_data, bytes);</div><div class="line"></div><div class="line">  <span class="comment">// capture my_kernel to do computation on gpu_data</span></div><div class="line">  <a class="code" href="classtf_1_1cudaTask.html">tf::cudaTask</a> kernel = capturer.<a class="code" href="classtf_1_1cudaFlowCapturerBase.html#adf651356def71f613c589c29588398c2">on</a>([&amp;](cudaStream_t stream){  </div><div class="line">    my_kernel&lt;&lt;&lt;grid, block, shm_size, stream&gt;&gt;&gt;(gpu_data, arg1, arg2, arg3, ...);</div><div class="line">  });</div><div class="line"></div><div class="line">  <span class="comment">// copy data from gpu_data to host_data</span></div><div class="line">  <a class="code" href="classtf_1_1cudaTask.html">tf::cudaTask</a> d2h = capturer.<a class="code" href="classtf_1_1cudaFlowCapturerBase.html#a38adf66dfcc0829708db653d153a83e2">memcpy</a>(host_data, gpu_data, bytes);</div><div class="line">  </div><div class="line">  h2d.<a class="code" href="classtf_1_1cudaTask.html#abdd68287ec4dff4216af34d1db44d1b4">precede</a>(kernel);</div><div class="line">  kernel.<a class="code" href="classtf_1_1cudaTask.html#abdd68287ec4dff4216af34d1db44d1b4">precede</a>(d2h);</div><div class="line">});</div></div><!-- fragment --><p>By inheriting <a class="el" href="classtf_1_1cudaFlowCapturerBase.html" title="base class to construct a CUDA task graph through stream capture ">tf::cudaFlowCapturerBase</a>, you can create your own capturer. The <a class="el" href="classtf_1_1cudaFlowCapturer.html" title="class for building a CUDA task dependency graph through stream capture ">tf::cudaFlowCapturer</a> has a factory interface, <a class="el" href="classtf_1_1cudaFlowCapturer.html#a49dd9727aeb5bccd7d826c6b11a6882d" title="creates a custom capturer derived from tf::cudaFlowCapturerBase ">tf::cudaFlowCapturer::make_capturer</a>, for users to create custom capturers with lifetimes managed by the capturer. This is convenient when you need certain objects alive during the capturing.</p>
<p>The following example shows a custom capturer on top of <a href="https://docs.nvidia.com/cuda/cublas/index.html">cuBLAS</a> that includes the <code>cublasHandle</code> object <code>_handle</code> to associate operations with streams.</p>
<div class="fragment"><div class="line"><span class="keyword">class </span>MyCapturer : <span class="keyword">public</span> <a class="code" href="classtf_1_1cudaFlowCapturerBase.html">tf::cudaFlowCapturerBase</a> {</div><div class="line"></div><div class="line">  <span class="keyword">public</span>:</div><div class="line"></div><div class="line">  MyCapturer(args_to_construct_MyCapturer) {</div><div class="line">    <span class="comment">// ...</span></div><div class="line">  }</div><div class="line"></div><div class="line">  <a class="code" href="classtf_1_1cudaTask.html">tf::cudaTask</a> matrix_multiplication(args...) {</div><div class="line">    <span class="keywordflow">return</span> on([<span class="keyword">this</span>, args...](cudaStream_t stream){</div><div class="line">      cublasSetStream(_handle, stream);</div><div class="line">      cublasSgemm(_handle, args...);</div><div class="line">    });</div><div class="line">  }</div><div class="line"></div><div class="line">  <span class="keyword">private</span>:</div><div class="line"></div><div class="line">    cuBlasHandle_t _handle;</div><div class="line">    Data _other_internal_data;</div><div class="line">};</div><div class="line"></div><div class="line"><a class="code" href="classtf_1_1Task.html">tf::Task</a> task = taskflow.<a class="code" href="classtf_1_1FlowBuilder.html#a60d7a666cab71ecfa3010b2efb0d6b57">emplace</a>([&amp;](<a class="code" href="classtf_1_1cudaFlowCapturer.html">tf::cudaFlowCapturer</a>&amp; capturer){</div><div class="line"></div><div class="line">  MyCapturer* mc = capturer.<a class="code" href="classtf_1_1cudaFlowCapturer.html#a49dd9727aeb5bccd7d826c6b11a6882d">make_capturer</a>&lt;MyCapturer&gt;(args_to_construct_MyCapturer);</div><div class="line"></div><div class="line">  <a class="code" href="classtf_1_1cudaTask.html">tf::cudaTask</a> task_1 = mc-&gt;matrix_multiplication(args...);</div><div class="line">  <a class="code" href="classtf_1_1cudaTask.html">tf::cudaTask</a> task_2 = capturer.<a class="code" href="classtf_1_1cudaFlowCapturerBase.html#adf651356def71f613c589c29588398c2">on</a>([&amp;](cudaStream_t stream){</div><div class="line">    other_kernel&lt;&lt;&lt;grid, block, shm_size, stream&gt;&gt;&gt;(other_args...);</div><div class="line">  });</div><div class="line"></div><div class="line">  task_1.<a class="code" href="classtf_1_1cudaTask.html#abdd68287ec4dff4216af34d1db44d1b4">precede</a>(task_2);</div><div class="line">});</div></div><!-- fragment --><h1><a class="anchor" id="C7_CaptureOnASpecificGPU"></a>
Capture on a Specific GPU</h1>
<p>You can capture a cudaFlow on a specific GPU by calling <a class="el" href="classtf_1_1FlowBuilder.html#afdf47fd1a358fb64f8c1b89e2a393169" title="creates a cudaflow task on the given device ">tf::Taskflow::emplace_on</a>. By default, a cudaFlow runs on the current GPU associated with the caller, which is typically 0. Similar to <a class="el" href="chapter6.html#C6_run_a_cudaflow_on_multiple_gpus">Run a cudaFlow on Multiple GPUs</a>, you can emplace a cudaFlowCapturer on a specific GPU.</p>
<div class="fragment"><div class="line"><a class="code" href="classtf_1_1Task.html">tf::Task</a> task = taskflow.<a class="code" href="classtf_1_1FlowBuilder.html#afdf47fd1a358fb64f8c1b89e2a393169">emplace_on</a>([](<a class="code" href="classtf_1_1cudaFlowCapturer.html">tf::cudaFlowCapturer</a>&amp; capturer){</div><div class="line">  <span class="comment">// here, capturer is under GPU device context 2</span></div><div class="line">  <span class="comment">// ...</span></div><div class="line">}, 2);</div></div><!-- fragment --><p>The above example creates a capturer on GPU 2. When the executor runs the callable, it switches to GPU 2 and all the functions within the callable are called under this context. Keep in mind, Taskflow does not deal with memory. It is users' responsibility to ensure the allocated memory to sit in the right context. In this example, the GPU memory you created must come from context 2 as well. You may use unified shared memory (i.e., <code>cudaMallocManaged</code>) to avoid this pitfall, given that the extra cost of automatic memory migration is not hurting your application performance. </p>
</div></div><!-- contents -->
</div><!-- doc-content -->
<!-- start footer part -->
<div id="nav-path" class="navpath"><!-- id is needed for treeview function! -->
  <ul>
    <li class="navelem"><a class="el" href="Cookbook.html">Cookbook</a></li>
    <li class="footer">Generated by
    <a href="http://www.doxygen.org/index.html">
    <img class="footer" src="doxygen.png" alt="doxygen"/></a> 1.8.14 </li>
  </ul>
</div>
</body>
</html>
